/*
编译方法：${ROCM_PATH}/llvm/bin/clang++ -o cl_MatrixTranspose_process -I${ROCM_PATH}/opencl/include -lamdocl64 -lpthread -lOpenCL cl_MatrixTranspose_process.cpp
*/
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <CL/cl.h>
#include <unistd.h>
#include <sys/types.h> 
#include <sys/wait.h>
#include <thread>

#define WIDTH 1024
#define NUM (WIDTH * WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4

// OpenCL内核代码
const char* kernelSource = 
"__kernel void matrixTranspose(__global float* out, __global float* in, const int width) {\n"
"    int x = get_global_id(0);\n"
"    int y = get_global_id(1);\n"
"    out[y * width + x] = in[x * width + y];\n"
"}\n";

// CPU参考实现
void matrixTransposeCPUReference(float* output, float* input, const unsigned int width) {
    for (unsigned int j = 0; j < width; j++) {
        for (unsigned int i = 0; i < width; i++) {
            output[i * width + j] = input[j * width + i];
        }
    }
}

void demo() {
    cl_int err;
    float* Matrix = (float*)malloc(NUM * sizeof(float));
    float* TransposeMatrix = (float*)malloc(NUM * sizeof(float));
    float* cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));

    // 初始化输入数据
    for (int i = 0; i < NUM; i++) {
        Matrix[i] = (float)i * 10.0f;
    }

    // 1. 获取OpenCL平台
    cl_platform_id platform;
    err = clGetPlatformIDs(1, &platform, NULL);
    
    // 2. 获取GPU设备
    cl_device_id device;
    err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    
    // 3. 创建上下文
    cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
    
    // 4. 创建命令队列
    cl_command_queue queue = clCreateCommandQueue(context, device, 0, &err);
    
    // 5. 创建内存缓冲区
    cl_mem gpuMatrix = clCreateBuffer(context, CL_MEM_READ_ONLY, NUM * sizeof(float), NULL, &err);
    cl_mem gpuTransposeMatrix = clCreateBuffer(context, CL_MEM_WRITE_ONLY, NUM * sizeof(float), NULL, &err);
    
    // 6. 传输数据到设备
    err = clEnqueueWriteBuffer(queue, gpuMatrix, CL_TRUE, 0, NUM * sizeof(float), Matrix, 0, NULL, NULL);
    
    // 7. 创建并构建程序
    cl_program program = clCreateProgramWithSource(context, 1, &kernelSource, NULL, &err);
    err = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
    
    // 8. 创建内核
    cl_kernel kernel = clCreateKernel(program, "matrixTranspose", &err);
    
    // 9. 设置内核参数
    int width_arg = WIDTH;
    err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &gpuTransposeMatrix);
    err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &gpuMatrix);
    err = clSetKernelArg(kernel, 2, sizeof(int), &width_arg);
    
    // 10. 执行内核
    size_t globalSize[2] = {WIDTH, WIDTH};
    size_t localSize[2] = {THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y};
    err = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, globalSize, localSize, 0, NULL, NULL);
    
    // 11. 读取结果
    err = clEnqueueReadBuffer(queue, gpuTransposeMatrix, CL_TRUE, 0, NUM * sizeof(float), TransposeMatrix, 0, NULL, NULL);
    
    // CPU参考实现
    matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
    
    // 验证结果
    int errors = 0;
    double eps = 1.0E-6;
    for (int i = 0; i < NUM; i++) {
        if (fabs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) {
            errors++;
        }
    }
    
    if (errors != 0) {
        printf("FAILED: %d errors\n", errors);
    } else {
        printf("PASSED!\n");
    }
    
    // 释放资源
    clReleaseMemObject(gpuMatrix);
    clReleaseMemObject(gpuTransposeMatrix);
    clReleaseKernel(kernel);
    clReleaseProgram(program);
    clReleaseCommandQueue(queue);
    clReleaseContext(context);
    
    free(Matrix);
    free(TransposeMatrix);
    free(cpuTransposeMatrix);
}

int main(int argc, char **argv)
{
    printf("> %s Starting...\n", argv[0]);
    pid_t pid = fork();
    fprintf(stdout, "parent pid = %u , ppid() = %u\n",getpid(), getppid());
    if (pid == 0)
    {//child process 1
        fprintf(stdout, "child1 pid = %u , ppid() = %u\n",getpid(), getppid());
	pid = fork();
	if (pid == 0)
	{//child process 2
        fprintf(stdout, "child2 pid = %u , ppid() = %u\n",getpid(), getppid());
	    std::thread testthread = std::thread(demo);
	    demo();
	    testthread.join();
	    return 0;
	}
    }
    demo();
    waitpid(pid, NULL, 0);
    return 0;
}